#include "cuda_lighting.h"
#include "cudaSetup.h"

#include "fs_SSFlags.h"

#include <stdio.h>

#ifdef _WIN32
#include <windows.h>
#endif

#include <GL/glew.h>

#define MAX_BLOCK_SIZE 768
#define MAX_POINT_LIGHTS 4096

bool _cudaInit = false;
bool _use41 = false;
unsigned int *numTiles = NULL;
int *shaderIDs = NULL;
int *d_shaderIDs = NULL;
dim3 lastBlock = dim3(0,0);
int lastWidth = 0;
int lastHeight = 0;

uchar4 *d_outData = NULL;

int count = 0;

dim3 m;

#define pointFactor .4f * .8f

//#define ALT_C_VERSION

void cul_deleteCreatedArrays() {
	delete numTiles;
	numTiles = NULL;
	delete shaderIDs;
	shaderIDs = NULL;
}

extern "C" {

texture<uchar4, cudaTextureType2D, cudaReadModeElementType> depthTexRef;
texture<uchar4, cudaTextureType2D, cudaReadModeElementType> normalTexRef;
texture<uchar4, cudaTextureType2D, cudaReadModeElementType> diffuseTexRef;
texture<uchar4, cudaTextureType2D, cudaReadModeElementType> miscTexRef;
texture<uchar4, cudaTextureType2D, cudaReadModeElementType> miscBTexRef;
texture<float4, cudaTextureType2D, cudaReadModeElementType> outTexRef;
texture<uchar4, cudaTextureType2D, cudaReadModeElementType> edgeTexRef;

__device__ float4 mMult(float *mat, float4 vec) {
	float4 out;	
	out.x = mat[0]*vec.x + mat[4]*vec.y + mat[8]*vec.z + mat[12]*vec.w;
	out.y = mat[1]*vec.x + mat[5]*vec.y + mat[9]*vec.z + mat[13]*vec.w;
	out.z = mat[2]*vec.x + mat[6]*vec.y + mat[10]*vec.z + mat[14]*vec.w;
	out.w = mat[3]*vec.x + mat[7]*vec.y + mat[11]*vec.z + mat[15]*vec.w;
	return out;
}

__device__ float3 mMult3(float *mat, float4 vec) {
	float3 out;
	out.x = mat[0]*vec.x + mat[4]*vec.y + mat[8]*vec.z + mat[12]*vec.w;
	out.y = mat[1]*vec.x + mat[5]*vec.y + mat[9]*vec.z + mat[13]*vec.w;
	out.z = mat[2]*vec.x + mat[6]*vec.y + mat[10]*vec.z + mat[14]*vec.w;

	return out;
}

// Assumes the w coordinate was 1.0
__device__ float3 modelPosition(float3 viewPos, float *invView) {
	float3 temp;
	temp.x = invView[0]*viewPos.x + invView[4]*viewPos.y + invView[8]*viewPos.z + invView[12]*1.0f;
	temp.y = invView[1]*viewPos.x + invView[5]*viewPos.y + invView[9]*viewPos.z + invView[13]*1.0f;
	temp.z = invView[2]*viewPos.x + invView[6]*viewPos.y + invView[10]*viewPos.z + invView[14]*1.0f;
	return temp;
}

__device__ float3 normalize(float3 in) {
	float l = sqrt(in.x*in.x + in.y*in.y + in.z*in.z);
	return make_float3(in.x/l, in.y/l, in.z/l);
}

__device__ float dist(float3 in1, float3 in2) {
	float3 sub = make_float3(in1.x - in2.x, in1.y - in2.y, in1.z - in2.z);
	return sqrt(sub.x*sub.x + sub.y*sub.y + sub.z*sub.z);
}

__device__ float3 vsPosition(float depth, float x, float y, float *mat) {
	float4 sPos = make_float4(2.0f*x - 1.0f, 2.0f*y - 1.0f, 2.0f*depth - 1.0f, 1.0f);
	sPos = mMult(mat, sPos);
	return make_float3(sPos.x/sPos.w, sPos.y/sPos.w, sPos.z/sPos.w);
}

__device__ float3 transformNormal(float3 normal, float *invTransposed) {
	float3 temp;
	temp.x = invTransposed[0]*normal.x + invTransposed[3]*normal.y + invTransposed[6]*normal.z;
	temp.y = invTransposed[1]*normal.x + invTransposed[4]*normal.y + invTransposed[7]*normal.z;
	temp.z = invTransposed[2]*normal.x + invTransposed[5]*normal.y + invTransposed[8]*normal.z;
	return temp;
}

__device__ float dMax(float a, float b) {
	return a > b ? a : b;
}

__device__ float dot(float3 a, float3 b) {
	return a.x*b.x + a.y*b.y + a.z*b.z;
}

__device__ float3 colorMult(float3 a, float3 b) {
	return make_float3(a.x*b.x, a.y*b.y, a.z*b.z);
}

__device__ int iMin(int a, int b) {
	return a < b ? a : b;
}

__device__ float lDepth(float d, float n, float f) {
	return (2.0f*n)/(f + n - d*(f - n));
}

#ifndef ALT_C_VERSION

__global__ void cudaCalculation(unsigned int numPointLights, float *plFloats, int *srBounds,
								  int width, int height, uchar4 *outData) {

	__shared__ float mat[48];
	__shared__ bool tileInRect[MAX_BLOCK_SIZE];
	__shared__ float pLightInfo[8];
	__shared__ bool needsProcessing;

	int tdx = threadIdx.x + blockDim.x*threadIdx.y;
	int blockS = blockDim.x * blockDim.y;
	if(tdx < warpSize)
		needsProcessing = false;

	// Load the inverse perspective and view matrices
	if(blockS < 32) {
		if(tdx < 16) {
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
			mat[tdx+16] = plFloats[MAX_POINT_LIGHTS*8 + 16 + tdx];
			mat[tdx+32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
		}
	}
	else if (blockS < 48) {
		if(tdx < 32)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
		if(tdx < 16)
			mat[tdx + 32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
	}
	else
		if(tdx < 48)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];

	// Calculate texture read coordinates
	ushort2 coords;
	coords.x = blockDim.x * blockIdx.x + threadIdx.x;
	coords.y = blockDim.y * blockIdx.y + threadIdx.y;
	bool onScreen = (coords.x < width) && (coords.y < height);

	// Get the depth and VS position
	uchar4 txRead = tex2D(depthTexRef, coords.x, coords.y);
	unsigned int depthUInt = txRead.x + txRead.y*256u + txRead.z*256u*256u + txRead.w*256u*256u*256u;
	if(depthUInt != 0)
		needsProcessing = true;
	__syncthreads(); //  First time that iP is used
	if(!needsProcessing) {
		if(onScreen)
			outData[coords.x + width*coords.y] = make_uchar4(0,0,0,255);
		return;
	}
	float depth = __int_as_float(depthUInt);
	float3 pos = vsPosition(depth, (float)coords.x/width, (float)coords.y/height, &mat[16]);
	pos = modelPosition(pos, &mat[0]);

	uchar4 diffuse;
	if(onScreen) {
		diffuse = tex2D(diffuseTexRef, coords.x, coords.y);
	}

	// Get the normal
	txRead = tex2D(normalTexRef, coords.x, coords.y);
	float n1 = txRead.x/(256.0f*255.0f) + txRead.y/255.0f;
	float n2 = txRead.z/(256.0f*255.0f) + txRead.w/255.0f;
	n1 = n1*4.0f - 2.0f;
	n2 = n2*4.0f - 2.0f;
	float f = n1*n1 + n2*n2;
	float g = sqrt(1.0f - .25f*f);
	float3 normal = make_float3(n1*g, n2*g, 1.0f - .5f*f);
	normal = normalize(transformNormal(normal, &mat[32]));

	// Get the tile rect-bounds
	short4 pBounds;
	pBounds.x = blockIdx.x * blockDim.x;
	pBounds.y = blockIdx.y * blockDim.y;
	pBounds.z = pBounds.x + blockDim.x - 1;
	pBounds.w = pBounds.y + blockDim.y - 1;
	
	// Initialize color accumulation buffer
	float3 color_accum = make_float3(0,0,0);

	unsigned int numLoops = numPointLights / blockS + (numPointLights % blockS == 0 ? 0 : 1);
	// Loop through for all of the lights

	// Try declaring variables outside of loop
	float d, attenuate, factor;
	float3 lpos, ldir;
	unsigned int count = 0;
	for(unsigned int i = 0; i < numLoops; ++i) {
		int index = i*blockS + tdx;

		// Check to see if a light is in bounds
		if(index < numPointLights) {
			tileInRect[tdx] = !((srBounds[index] >= pBounds.z) ||
							    (srBounds[index + numPointLights] >= pBounds.w) ||
							    (srBounds[index + 2*numPointLights] <= pBounds.x) ||
							    (srBounds[index + 3*numPointLights] <= pBounds.y));
		}
		else
			tileInRect[tdx] = false;

		__syncthreads();

		// For each light that was in bounds, calculate the light contribution
		for(unsigned int j = 0; j < blockS; ++j) {
			if(tileInRect[j]) {
				++count;
				// Load the info for a single light into shared memory
				__syncthreads();
				if(tdx < 8)
						pLightInfo[tdx] = plFloats[(i*blockS+j)*8 + tdx];

				__syncthreads();
				// Only perform light calculation if the thread corresponds to an on-screen coordinate
				if(onScreen) {
					lpos = make_float3(pLightInfo[0], pLightInfo[1], pLightInfo[2]);

					d = dist(lpos, pos);

					attenuate = 1.0f;
					if(d > pLightInfo[6])
						attenuate = .0f;
					else if(d > pLightInfo[7])
						attenuate -= (d - pLightInfo[7])/(pLightInfo[6] - pLightInfo[7]);
				
					// N*L*attenuation
					ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
					factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

					// Accumulate the color contribution from the current light
					color_accum.x += factor*pLightInfo[3];
					color_accum.y += factor*pLightInfo[4];
					color_accum.z += factor*pLightInfo[5];
				}
			}
		}
		__syncthreads();
	}

	// Multiply light accumulation by the texture color
	if(onScreen)
		//outData[coords.x + width*coords.y] = make_uchar4(count*64, count*16, count*2, 1.0f);
		outData[coords.x + width*coords.y] = make_uchar4(pointFactor*color_accum.x*diffuse.x,
														 pointFactor*color_accum.y*diffuse.y,
														 pointFactor*color_accum.z*diffuse.z, 255);
}

#else

__global__ void cudaCalculation(unsigned int numPointLights, float *plFloats, int *srBounds,
								int width, int height, uchar4 *outData) {

	__shared__ float mat[48];
	__shared__ bool tileInRect[MAX_BLOCK_SIZE];
	__shared__ float pLightInfo[8*2];

	int tdx = threadIdx.x + blockDim.x*threadIdx.y;
	int blockS = blockDim.x * blockDim.y;

	// Load the inverse perspective, inverse view, and inverse view transposed matrices
	if(blockS < 32) {
		if(tdx < 16) {
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
			mat[tdx+16] = plFloats[MAX_POINT_LIGHTS*8 + 16 + tdx];
			mat[tdx+32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
		}
	}
	else if (blockS < 48) {
		if(tdx < 32)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
		if(tdx < 16)
			mat[tdx + 32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
	}
	else
		if(tdx < 48)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];


	// Calculate texture read coordinates
	ushort2 coords;
	coords.x = blockDim.x * blockIdx.x + threadIdx.x;
	coords.y = blockDim.y * blockIdx.y + threadIdx.y;
	bool onScreen = (coords.x < width) && (coords.y < height);

	// Get the normal
	uchar4 txRead = tex2D(normalTexRef, coords.x, coords.y);
	float n1 = txRead.x/(256.0f*255.0f) + txRead.y/255.0f;
	float n2 = txRead.z/(256.0f*255.0f) + txRead.w/255.0f;
	n1 = n1*4.0f - 2.0f;
	n2 = n2*4.0f - 2.0f;
	float f = n1*n1 + n2*n2;
	float g = sqrt(1.0f - .25f*f);
	float3 normal = make_float3(n1*g, n2*g, 1.0f - .5f*f);

	// Get the depth and VS position
	txRead = tex2D(depthTexRef, coords.x, coords.y);
	float depth = txRead.x/(256.0f*256.0f*256.0f*255.0f) +
				  txRead.y/(256.0f*256.0f*255.0f) +
				  txRead.z/(256.0f*255.0f) +
				  txRead.w/(255.0f);

	__syncthreads(); //  First time that loaded matrix values are used
	float3 pos = vsPosition(depth, (float)coords.x/width, (float)coords.y/height, &mat[16]);
	pos = modelPosition(pos, &mat[0]);
	normal = normalize(transformNormal(normal, &mat[32]));

	uchar4 diffuse;
	if(onScreen) {
		diffuse = tex2D(diffuseTexRef, coords.x, coords.y);
	}

	// Get the tile rect-bounds
	ushort4 pBounds;
	pBounds.x = blockIdx.x * blockDim.x;
	pBounds.y = blockIdx.y * blockDim.y;
	pBounds.z = pBounds.x + blockDim.x - 1;
	pBounds.w = pBounds.y + blockDim.y - 1;
	
	// Initialize color accumulation buffer
	float3 color_accum = make_float3(0,0,0);

	unsigned int numLoops = numPointLights / blockS + (numPointLights % blockS == 0 ? 0 : 1);
	unsigned int loader = 0;

	// Try declaring variables outside of loop
	float d, attenuate, factor;
	float3 lpos, ldir;

	// Loop through for all of the lights
	for(unsigned int i = 0; i < numLoops; ++i) {
		int index = i*blockS + tdx;

		// Check to see if a light is in bounds
		if(index < numPointLights) {
			tileInRect[tdx] = !((srBounds[index] >= pBounds.z) ||
							    (srBounds[index + numPointLights] >= pBounds.w) ||
							    (srBounds[index + 2*numPointLights] <= pBounds.x) ||
							    (srBounds[index + 3*numPointLights] <= pBounds.y));
		}
		else
			tileInRect[tdx] = false;

		__syncthreads();

		// For each light that was in bounds, calculate the light contribution
		for(unsigned int j = 0; j < blockS; ++j) {
			if(tileInRect[j]) {
				__syncthreads();
				if(loader == 0) {
					// Load the info for a single light into shared memory
					if(tdx < 8) {
						pLightInfo[tdx] = plFloats[(i*blockS+j)*8 + tdx];
					}
					loader = 1;
				}
				else if(loader == 1) {
					// Load the info for a single light into shared memory
					if(blockS > 47)
						if(tdx > 31 && tdx < 40)
							pLightInfo[tdx - 32 + 8] = plFloats[(i*blockS+j)*8 + tdx - 32];
					else
						if(tdx < 8)
							pLightInfo[tdx + 8] = plFloats[(i*blockS+j)*8 + tdx];

					// Only perform light calculation if the thread corresponds to an on-screen coordinate
					if(onScreen) {
						lpos = make_float3(pLightInfo[0], pLightInfo[1], pLightInfo[2]);

						d = dist(lpos, pos);

						attenuate = 1.0f;
						if(d > pLightInfo[6])
							attenuate = .0f;
						else if(d > pLightInfo[7])
							attenuate -= (d - pLightInfo[7])/(pLightInfo[6] - pLightInfo[7]);
					
						// N*L*attenuation
						ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
						factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

						// Accumulate the color contribution from the current light
						color_accum.x += factor*pLightInfo[3];
						color_accum.y += factor*pLightInfo[4];
						color_accum.z += factor*pLightInfo[5];
					}
					loader = 2;
				}
				else if(loader == 2) {
					if(tdx < 8)
						pLightInfo[tdx] = plFloats[(i*blockS+j)*8 + tdx];

					// Only perform light calculation if the thread corresponds to an on-screen coordinate
					if(onScreen) {
						lpos = make_float3(pLightInfo[0+8], pLightInfo[1+8], pLightInfo[2+8]);

						d = dist(lpos, pos);

						attenuate = 1.0f;
						if(d > pLightInfo[6+8])
							attenuate = .0f;
						else if(d > pLightInfo[7+8])
							attenuate -= (d - pLightInfo[7+8])/(pLightInfo[6+8] - pLightInfo[7+8]);
					
						// N*L*attenuation
						ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
						factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

						// Accumulate the color contribution from the current light
						color_accum.x += factor*pLightInfo[3+8];
						color_accum.y += factor*pLightInfo[4+8];
						color_accum.z += factor*pLightInfo[5+8];
					}
					loader = 1;
				}
			}
				
				
		}
		__syncthreads();
		if(loader == 1) {
			// Only perform light calculation if the thread corresponds to an on-screen coordinate
			if(onScreen) {
				lpos = make_float3(pLightInfo[0], pLightInfo[1], pLightInfo[2]);

				d = dist(lpos, pos);

				attenuate = 1.0f;
				if(d > pLightInfo[6])
					attenuate = .0f;
				else if(d > pLightInfo[7])
					attenuate -= (d - pLightInfo[7])/(pLightInfo[6] - pLightInfo[7]);
				
				// N*L*attenuation
				ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
				factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

				// Accumulate the color contribution from the current light
				color_accum.x += factor*pLightInfo[3];
				color_accum.y += factor*pLightInfo[4];
				color_accum.z += factor*pLightInfo[5];
			}
		}
		else if(loader == 2) {
			// Only perform light calculation if the thread corresponds to an on-screen coordinate
			if(onScreen) {
				lpos = make_float3(pLightInfo[0+8], pLightInfo[1+8], pLightInfo[2+8]);

				d = dist(lpos, pos);

				attenuate = 1.0f;
				if(d > pLightInfo[6+8])
					attenuate = .0f;
				else if(d > pLightInfo[7+8])
					attenuate -= (d - pLightInfo[7+8])/(pLightInfo[6+8] - pLightInfo[7+8]);
				
				// N*L*attenuation
				ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
				factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

				// Accumulate the color contribution from the current light
				color_accum.x += factor*pLightInfo[3+8];
				color_accum.y += factor*pLightInfo[4+8];
				color_accum.z += factor*pLightInfo[5+8];
			}
		}
		__syncthreads();
		loader = 0;
	}

	// Multiply light accumulation by the texture color
	if(onScreen)
			outData[coords.x + width*coords.y] = make_uchar4(iMin(int(pointFactor*color_accum.x*diffuse.x), 255),
															 iMin(int(pointFactor*color_accum.y*diffuse.y), 255),
															 iMin(int(pointFactor*color_accum.z*diffuse.z), 255), 255);	
}
#endif // #ifndef ALT_C_VERSION

#ifdef ALT_C_VERSION

__global__ void cudaCalculation41(unsigned int numPointLights, float *plFloats, int *srBounds,
								  int width, int height, uchar4 *outData) {

	__shared__ float mat[48];
	__shared__ bool tileInRect[MAX_BLOCK_SIZE];
	__shared__ float pLightInfo[8*2];
	__shared__ bool needsProcessing;

	int tdx = threadIdx.x + blockDim.x*threadIdx.y;
	int blockS = blockDim.x * blockDim.y;
	if(tdx < warpSize)
		needsProcessing = false;

	// Load the inverse perspective, inverse view, and inverse view transposed matrices
	if(blockS < 32) {
		if(tdx < 16) {
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
			mat[tdx+16] = plFloats[MAX_POINT_LIGHTS*8 + 16 + tdx];
			mat[tdx+32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
		}
	}
	else if (blockS < 48) {
		if(tdx < 32)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
		if(tdx < 16)
			mat[tdx + 32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
	}
	else
		if(tdx < 48)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];

	// Calculate texture read coordinates
	ushort2 coords;
	coords.x = blockDim.x * blockIdx.x + threadIdx.x;
	coords.y = blockDim.y * blockIdx.y + threadIdx.y;
	bool onScreen = (coords.x < width) && (coords.y < height);

	// Get the depth and VS position
	uchar4 txRead = tex2D(depthTexRef, coords.x, coords.y);
	unsigned int depthUInt = txRead.x + txRead.y*256u + txRead.z*256u*256u + txRead.w*256u*256u*256u;
	if(depthUInt != 0)
		needsProcessing = true;
	__syncthreads(); //  First time that iP is used
	if(!needsProcessing) {
		if(onScreen)
			outData[coords.x + width*coords.y] = make_uchar4(0,0,0,255);
		return;
	}
	float depth = __int_as_float(depthUInt);
	float3 pos = vsPosition(depth, (float)coords.x/width, (float)coords.y/height, &mat[16]);
	pos = modelPosition(pos, &mat[0]);

	uchar4 diffuse;
	if(onScreen) {
		diffuse = tex2D(diffuseTexRef, coords.x, coords.y);
	}

	// Get the normal
	txRead = tex2D(normalTexRef, coords.x, coords.y);
	float n1 = txRead.x/(256.0f*255.0f) + txRead.y/255.0f;
	float n2 = txRead.z/(256.0f*255.0f) + txRead.w/255.0f;
	n1 = n1*4.0f - 2.0f;
	n2 = n2*4.0f - 2.0f;
	float f = n1*n1 + n2*n2;
	float g = sqrt(1.0f - .25f*f);
	float3 normal = make_float3(n1*g, n2*g, 1.0f - .5f*f);
	normal = normalize(transformNormal(normal, &mat[32]));

	// Get the tile rect-bounds
	short4 pBounds;
	pBounds.x = blockIdx.x * blockDim.x;
	pBounds.y = blockIdx.y * blockDim.y;
	pBounds.z = pBounds.x + blockDim.x - 1;
	pBounds.w = pBounds.y + blockDim.y - 1;
	
	// Initialize color accumulation buffer
	float3 color_accum = make_float3(0,0,0);

	unsigned int numLoops = numPointLights / blockS + (numPointLights % blockS == 0 ? 0 : 1);
	unsigned int loader = 0;

	// Try declaring variables outside of loop
	float d, attenuate, factor;
	float3 lpos, ldir;

	// Loop through for all of the lights
	for(unsigned int i = 0; i < numLoops; ++i) {
		int index = i*blockS + tdx;

		// Check to see if a light is in bounds
		if(index < numPointLights) {
			tileInRect[tdx] = !((srBounds[index] >= pBounds.z) ||
							    (srBounds[index + numPointLights] >= pBounds.w) ||
							    (srBounds[index + 2*numPointLights] <= pBounds.x) ||
							    (srBounds[index + 3*numPointLights] <= pBounds.y));
		}
		else
			tileInRect[tdx] = false;

		__syncthreads();

		// For each light that was in bounds, calculate the light contribution
		for(unsigned int j = 0; j < blockS; ++j) {
			if(tileInRect[j]) {
				__syncthreads();
				if(loader == 0) {
					// Load the info for a single light into shared memory
					if(tdx < 8) {
						pLightInfo[tdx] = plFloats[(i*blockS+j)*8 + tdx];
					}
					loader = 1;
				}
				else if(loader == 1) {
					// Load the info for a single light into shared memory
					if(blockS > 47)
						if(tdx > 31 && tdx < 40)
							pLightInfo[tdx - 32 + 8] = plFloats[(i*blockS+j)*8 + tdx - 32];
					else
						if(tdx < 8)
							pLightInfo[tdx + 8] = plFloats[(i*blockS+j)*8 + tdx];

					// Only perform light calculation if the thread corresponds to an on-screen coordinate
					if(onScreen) {
						lpos = make_float3(pLightInfo[0], pLightInfo[1], pLightInfo[2]);

						d = dist(lpos, pos);

						attenuate = 1.0f;
						if(d > pLightInfo[6])
							attenuate = .0f;
						else if(d > pLightInfo[7])
							attenuate -= (d - pLightInfo[7])/(pLightInfo[6] - pLightInfo[7]);
					
						// N*L*attenuation
						ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
						factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

						// Accumulate the color contribution from the current light
						color_accum.x += factor*pLightInfo[3];
						color_accum.y += factor*pLightInfo[4];
						color_accum.z += factor*pLightInfo[5];
					}
					loader = 2;
				}
				else if(loader == 2) {
					if(tdx < 8)
						pLightInfo[tdx] = plFloats[(i*blockS+j)*8 + tdx];

					// Only perform light calculation if the thread corresponds to an on-screen coordinate
					if(onScreen) {
						lpos = make_float3(pLightInfo[0+8], pLightInfo[1+8], pLightInfo[2+8]);

						d = dist(lpos, pos);

						attenuate = 1.0f;
						if(d > pLightInfo[6+8])
							attenuate = .0f;
						else if(d > pLightInfo[7+8])
							attenuate -= (d - pLightInfo[7+8])/(pLightInfo[6+8] - pLightInfo[7+8]);
					
						// N*L*attenuation
						ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
						factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

						// Accumulate the color contribution from the current light
						color_accum.x += factor*pLightInfo[3+8];
						color_accum.y += factor*pLightInfo[4+8];
						color_accum.z += factor*pLightInfo[5+8];
					}
					loader = 1;
				}
			}
				
				
		}

		if(loader == 1) {
			// Only perform light calculation if the thread corresponds to an on-screen coordinate
			if(onScreen) {
				lpos = make_float3(pLightInfo[0], pLightInfo[1], pLightInfo[2]);

				d = dist(lpos, pos);

				attenuate = 1.0f;
				if(d > pLightInfo[6])
					attenuate = .0f;
				else if(d > pLightInfo[7])
					attenuate -= (d - pLightInfo[7])/(pLightInfo[6] - pLightInfo[7]);
				
				// N*L*attenuation
				ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
				factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

				// Accumulate the color contribution from the current light
				color_accum.x += factor*pLightInfo[3];
				color_accum.y += factor*pLightInfo[4];
				color_accum.z += factor*pLightInfo[5];
			}
		}
		else if(loader == 2) {
			// Only perform light calculation if the thread corresponds to an on-screen coordinate
			if(onScreen) {
				lpos = make_float3(pLightInfo[0+8], pLightInfo[1+8], pLightInfo[2+8]);

				d = dist(lpos, pos);

				attenuate = 1.0f;
				if(d > pLightInfo[6+8])
					attenuate = .0f;
				else if(d > pLightInfo[7+8])
					attenuate -= (d - pLightInfo[7+8])/(pLightInfo[6+8] - pLightInfo[7+8]);
				
				// N*L*attenuation
				ldir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
				factor = dMax(dot(ldir,normal), 0.0f)*attenuate;

				// Accumulate the color contribution from the current light
				color_accum.x += factor*pLightInfo[3+8];
				color_accum.y += factor*pLightInfo[4+8];
				color_accum.z += factor*pLightInfo[5+8];
			}
		}
		__syncthreads();
		loader = 0;
	}

	// Multiply light accumulation by the texture color
	if(onScreen)
			outData[coords.x + width*coords.y] = make_uchar4(iMin(int(pointFactor*color_accum.x*diffuse.x), 255),
															 iMin(int(pointFactor*color_accum.y*diffuse.y), 255),
															 iMin(int(pointFactor*color_accum.z*diffuse.z), 255), 255);	
}
#else
__global__ void cudaCalculation41(unsigned int numPointLights, float *plFloats, int *srBounds,
								  int width, int height, uchar4 *outData) {

	__shared__ float mat[48];
	__shared__ bool tileInRect[MAX_BLOCK_SIZE];
	__shared__ float pLightInfo[8];
	__shared__ bool needsProcessing;

	int tdx = threadIdx.x + blockDim.x*threadIdx.y;
	int blockS = blockDim.x * blockDim.y;
	if(tdx < warpSize)
		needsProcessing = false;

	// Load the inverse perspective, inverse view, and inverse view transposed matrices
	if(blockS < 32) {
		if(tdx < 16) {
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
			mat[tdx+16] = plFloats[MAX_POINT_LIGHTS*8 + 16 + tdx];
			mat[tdx+32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
		}
	}
	else if (blockS < 48) {
		if(tdx < 32)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];
		if(tdx < 16)
			mat[tdx + 32] = plFloats[MAX_POINT_LIGHTS*8 + 32 + tdx];
	}
	else
		if(tdx < 48)
			mat[tdx] = plFloats[MAX_POINT_LIGHTS*8 + tdx];

	// Calculate texture read coordinates
	ushort2 coords;
	coords.x = blockDim.x * blockIdx.x + threadIdx.x;
	coords.y = blockDim.y * blockIdx.y + threadIdx.y;
	bool onScreen = (coords.x < width) && (coords.y < height);

	// Get the depth and VS position
	uchar4 txRead = tex2D(depthTexRef, coords.x, coords.y);
	unsigned int depthUInt = txRead.x + txRead.y*256u + txRead.z*256u*256u + txRead.w*256u*256u*256u;
	if(depthUInt != 0)
		needsProcessing = true;
	__syncthreads(); //  First time that iP is used
	if(!needsProcessing) {
		if(onScreen)
			outData[coords.x + width*coords.y] = make_uchar4(0,0,0,255);
		return;
	}
	float depth = __int_as_float(depthUInt);
	float3 pos = vsPosition(depth, (float)coords.x/width, (float)coords.y/height, &mat[16]);
	pos = modelPosition(pos, &mat[0]);

	uchar4 diffuse;
	if(onScreen) {
		diffuse = tex2D(diffuseTexRef, coords.x, coords.y);
	}

	// Get the normal
	txRead = tex2D(normalTexRef, coords.x, coords.y);
	float n1 = txRead.x/(256.0f*255.0f) + txRead.y/255.0f;
	float n2 = txRead.z/(256.0f*255.0f) + txRead.w/255.0f;
	n1 = n1*4.0f - 2.0f;
	n2 = n2*4.0f - 2.0f;
	float f = n1*n1 + n2*n2;
	float g = sqrt(1.0f - .25f*f);
	float3 normal = make_float3(n1*g, n2*g, 1.0f - .5f*f);
	normal = normalize(transformNormal(normal, &mat[32]));

	// Get the tile rect-bounds
	short4 pBounds;
	pBounds.x = blockIdx.x * blockDim.x;
	pBounds.y = blockIdx.y * blockDim.y;
	pBounds.z = pBounds.x + blockDim.x - 1;
	pBounds.w = pBounds.y + blockDim.y - 1;
	
	// Initialize color accumulation buffer
	float3 color_accum = make_float3(0.0f,0.0f,0.0f);

	unsigned int numLoops = numPointLights / blockS + (numPointLights % blockS == 0 ? 0 : 1);
	// Loop through for all of the lights
	for(unsigned int i = 0u; i < numLoops; ++i) {
		int index = i*blockS + tdx;

		// Check to see if a light is in bounds
		if(index < numPointLights) {
			tileInRect[tdx] = !((srBounds[index] >= pBounds.z) ||
							    (srBounds[index + numPointLights] >= pBounds.w) ||
							    (srBounds[index + 2*numPointLights] <= pBounds.x) ||
							    (srBounds[index + 3*numPointLights] <= pBounds.y));
		}
		else
			tileInRect[tdx] = false;

		__syncthreads();

		// For each light that was in bounds, calculate the light contribution
		for(unsigned int j = 0; j < blockS; ++j) {
			
			if(tileInRect[j]) {
				__syncthreads();
				// Load the info for a single light into shared memory

				if(tdx < 8) {
					pLightInfo[tdx] = plFloats[(i*blockS+j)*8 + tdx];
				}
				__syncthreads();

				// Only perform light calculation if the thread corresponds to an on-screen coordinate
				if(onScreen) {
					float3 lpos = make_float3(pLightInfo[0], pLightInfo[1], pLightInfo[2]);

					float d = dist(lpos, pos);

					float attenuate = 1.0f;
					if(d > pLightInfo[6])
						attenuate = .0f;
					else if(d > pLightInfo[7])
						attenuate -= (d - pLightInfo[7])/(pLightInfo[6] - pLightInfo[7]);
				
					// N*L*attenuation
					float3 lDir = normalize(make_float3(lpos.x - pos.x, lpos.y - pos.y, lpos.z - pos.z));
					float factor = dMax(dot(lDir,normal), 0.0f)*attenuate;

					// Accumulate the color contribution from the current light
					color_accum.x += factor*pLightInfo[3];
					color_accum.y += factor*pLightInfo[4];
					color_accum.z += factor*pLightInfo[5];
				}
			}
		}
		__syncthreads();
	}

	// Multiply light accumulation by the texture color
	if(onScreen)
			outData[coords.x + width*coords.y] = make_uchar4(iMin(int(pointFactor*color_accum.x*diffuse.x), 255),
															 iMin(int(pointFactor*color_accum.y*diffuse.y), 255),
															 iMin(int(pointFactor*color_accum.z*diffuse.z), 255), 255);	
}
#endif // #ifdef ALT_C_VERSION

__global__ void generateShaderIDs(int *shaderIDs, int width, int height, float n, float f) {

	__shared__ bool flagExists[5];
	__shared__ bool depthMinMax[6];

	int tdx = threadIdx.x + blockDim.x*threadIdx.y;
	if(tdx < warpSize) {
		flagExists[0] = false;
		flagExists[1] = false;
		flagExists[2] = false;
		flagExists[3] = false;
		flagExists[4] = false;
		depthMinMax[0] = false;
		depthMinMax[1] = false;
		depthMinMax[2] = false;
		depthMinMax[3] = false;
		depthMinMax[4] = false;
		depthMinMax[5] = false;
	}

	// Calculate texture read coordinates
	ushort2 coords;
	coords.x = blockDim.x * blockIdx.x + threadIdx.x;
	coords.y = blockDim.y * blockIdx.y + threadIdx.y;
	bool onScreen = (coords.x < width) && (coords.y < height);

	if(onScreen) {
		uchar4 miscRead = tex2D(miscTexRef, coords.x, coords.y);

		if((miscRead.w & SS_FLAG_NOT_SKY) != 0)
			flagExists[0] = true;

		if((miscRead.w & SS_FLAG_ANIMATED) != 0 || (tex2D(miscBTexRef, coords.x, coords.y).w & SS_FLAG_ANIMATED) != 0)
			flagExists[1] = true;

		if(tex2D(edgeTexRef, coords.x, coords.y).x != 0)
			flagExists[2] = true;

		uchar4 txRead = tex2D(depthTexRef, coords.x, coords.y);
		float depth = txRead.x/(256.0f*256.0f*256.0f*255.0f) +
					  txRead.y/(256.0f*256.0f*255.0f) +
					  txRead.z/(256.0f*255.0f) +
					  txRead.w/(255.0f);
		if(depth < .000001f)
			depth = 1.0f;

		depth = lDepth(depth, n, f);

		if(depth < .2f)
			depthMinMax[0] = true;
		else
			depthMinMax[2] = true;

		if(depth > .55f)
			depthMinMax[1] = true;
		else
			depthMinMax[3] = true;

		if(depth < .25f)
			depthMinMax[4] = true;

		if(depth > .45f)
			depthMinMax[5] = true;
	}
	__syncthreads();

	// Only single thread can process here
	if(tdx == 0) {

		if(!depthMinMax[4] && !depthMinMax[5]) {
		
		}
		else if(depthMinMax[2] && depthMinMax[3]) {
			flagExists[3] = true;
			flagExists[4] = false;
		}
		else if(depthMinMax[0] && !depthMinMax[2] && flagExists[0])
			flagExists[4] = true;
		else {
			flagExists[3] = true;
			flagExists[4] = true;
		}


		shaderIDs[blockIdx.x + gridDim.x*blockIdx.y] =
			SS_FLAG_NOT_SKY*(int)flagExists[0] +
			SS_FLAG_ANIMATED*(int)flagExists[1] +
			SS_FLAG_EDGE*(int)flagExists[2] +
			SS_FLAG_DEPTH*(int)flagExists[3] +
			SS_FLAG_MAX_DEPTH*(int)flagExists[4];
	}
}

__global__ void generateShaderIDs41(int *shaderIDs, int width, int height, float n, float f) {

	__shared__ bool flagExists[5];
	__shared__ bool depthMinMax[6];

	int tdx = threadIdx.x + blockDim.x*threadIdx.y;
	if(tdx < warpSize) {
		flagExists[0] = false;
		flagExists[1] = false;
		flagExists[2] = false;
		flagExists[3] = false;
		flagExists[4] = false;
		depthMinMax[0] = false;
		depthMinMax[1] = false;
		depthMinMax[2] = false;
		depthMinMax[3] = false;
		depthMinMax[4] = false;
		depthMinMax[5] = false;
	}

	// Calculate texture read coordinates
	ushort2 coords;
	coords.x = blockDim.x * blockIdx.x + threadIdx.x;
	coords.y = blockDim.y * blockIdx.y + threadIdx.y;
	bool onScreen = (coords.x < width) && (coords.y < height);

	if(onScreen) {
		uchar4 miscRead = tex2D(miscTexRef, coords.x, coords.y);

		if(miscRead.w > 0)
			flagExists[0] = true;

		if((miscRead.w & SS_FLAG_ANIMATED) != 0 || (tex2D(miscBTexRef, coords.x, coords.y).w & SS_FLAG_ANIMATED) != 0)
			flagExists[1] = true;

		if(tex2D(edgeTexRef, coords.x, coords.y).x > 100)
			flagExists[2] = true;

		uchar4 txRead = tex2D(depthTexRef, coords.x, coords.y);
		unsigned int depthUInt = txRead.x + txRead.y*256u + txRead.z*256u*256u + txRead.w*256u*256u*256u;
		float depth = __int_as_float(depthUInt);
		if(depth < .000001f)
			depth = 1.0f;

		depth = lDepth(depth, n, f);

		if(depth < .2f)
			depthMinMax[0] = true;
		else
			depthMinMax[2] = true;

		if(depth > .55f)
			depthMinMax[1] = true;
		else
			depthMinMax[3] = true;

		if(depth < .25f)
			depthMinMax[4] = true;

		if(depth > .45f)
			depthMinMax[5] = true;
	}
	__syncthreads();

	// Only single thread can process here
	if(tdx == 0) {

		if(!depthMinMax[4] && !depthMinMax[5]) {
		
		}
		else if(depthMinMax[2] && depthMinMax[3]) {
			flagExists[3] = true;
			flagExists[4] = false;
		}
		else if(depthMinMax[0] && !depthMinMax[2] && flagExists[0])
			flagExists[4] = true;
		else {
			flagExists[3] = true;
			flagExists[4] = true;
		}


		shaderIDs[blockIdx.x + gridDim.x*blockIdx.y] =
			SS_FLAG_NOT_SKY*(int)flagExists[0] +
			SS_FLAG_ANIMATED*(int)flagExists[1] +
			SS_FLAG_EDGE*(int)flagExists[2] +
			SS_FLAG_DEPTH*(int)flagExists[3] +
			SS_FLAG_MAX_DEPTH*(int)flagExists[4];
	}
}

} // closing for "extern C"

void setCUDAInit(bool init) {
	cudaChannelFormatDesc cfc;
	cfc.x = cfc.y = cfc.z = cfc.w = 8;
	cfc.f = cudaChannelFormatKindUnsigned;

	cudaChannelFormatDesc cfc_o;
	cfc_o.x = cfc_o.y = cfc_o.z = cfc_o.w = 16;
	cfc.f = cudaChannelFormatKindFloat;

	// Just stuck these settings here for convenience
	depthTexRef.addressMode[0] = cudaAddressModeClamp;
	depthTexRef.addressMode[1] = cudaAddressModeClamp;
	depthTexRef.filterMode = cudaFilterModePoint;
	depthTexRef.normalized = 0;
	depthTexRef.channelDesc = cfc;

	normalTexRef.addressMode[0] = cudaAddressModeClamp;
	normalTexRef.addressMode[1] = cudaAddressModeClamp;
	normalTexRef.filterMode = cudaFilterModePoint;
	normalTexRef.normalized = 0;
	normalTexRef.channelDesc = cfc;

	diffuseTexRef.addressMode[0] = cudaAddressModeClamp;
	diffuseTexRef.addressMode[1] = cudaAddressModeClamp;
	diffuseTexRef.filterMode = cudaFilterModePoint;
	diffuseTexRef.normalized = 0;
	diffuseTexRef.channelDesc = cfc;

	miscTexRef.addressMode[0] = cudaAddressModeClamp;
	miscTexRef.addressMode[1] = cudaAddressModeClamp;
	miscTexRef.filterMode = cudaFilterModePoint;
	miscTexRef.normalized = 0;
	miscTexRef.channelDesc = cfc;

	miscBTexRef.addressMode[0] = cudaAddressModeClamp;
	miscBTexRef.addressMode[1] = cudaAddressModeClamp;
	miscBTexRef.filterMode = cudaFilterModePoint;
	miscBTexRef.normalized = 0;
	miscBTexRef.channelDesc = cfc;

	outTexRef.addressMode[0] = cudaAddressModeClamp;
	outTexRef.addressMode[1] = cudaAddressModeClamp;
	outTexRef.filterMode = cudaFilterModePoint;
	outTexRef.normalized = 0;
	outTexRef.channelDesc = cfc_o;

	edgeTexRef.addressMode[0] = cudaAddressModeClamp;
	edgeTexRef.addressMode[1] = cudaAddressModeClamp;
	edgeTexRef.filterMode = cudaFilterModePoint;
	edgeTexRef.normalized = 0;
	edgeTexRef.channelDesc = cfc;

	_cudaInit = init;

	const char *result = (const char*)glGetString(GL_SHADING_LANGUAGE_VERSION);
	if(result)
		_use41 = (result[0] >= '4' && result[2] >= '1');
}

// Ignoring CUDA errors
void performCUDALighting(dim3 block, int width, int height, ds_point_light_device pld, bool aFrame) {
	if(!_cudaInit) {
		printf("The textures in cuda have not been registered. Cannot perform CUDA calculations.\n");
		return;
	}
	else if((block.x * block.y) % 2 != 0) {
		printf("Threads per block must be a multiple of 2.\n");
		return;
	}
	else if((block.x * block.y) < 16) {
		printf("Must have at least 16 threads per block.\n");
		return;
	}

	dim3 grid = dim3(width / block.x, height / block.y);
	if(width % block.x != 0)
		grid.x = grid.x + 1;
	if(height % block.y != 0)
		grid.y = grid.y + 1;

	if(!(lastBlock.x == block.x && lastBlock.y == block.y && width == lastWidth && height == lastHeight)) {
		delete numTiles;
		numTiles = NULL;
		delete shaderIDs;
		shaderIDs = NULL;
		cudaFree(d_shaderIDs);
		d_shaderIDs = NULL;
		lastBlock = block;
		lastWidth = width;
		lastHeight = height;

		cudaFree(d_outData);
		d_outData = NULL;
	}
	if(!d_outData) {
		cudaMalloc(&d_outData, sizeof(uchar4)*width*height);
	}

	// Get the pointers to the CUDA graphics resources
	cuda_resource_container crc = getCudaResources(aFrame);
	//cudaGraphicsResource_t resources[5] = {crc.depth, crc.diffuse, crc.misc, crc.normal, crc.out};

	// Map the textures to CUDA
	//cudaGraphicsMapResources(5, &(resources[0]));
	cudaGraphicsMapResources(1, &(crc.depth));
	cudaGraphicsMapResources(1, &(crc.diffuse));
	cudaGraphicsMapResources(1, &(crc.misc));
	cudaGraphicsMapResources(1, &(crc.normal));
	cudaGraphicsMapResources(1, &(crc.out));

	// Get device pointers to the textures
	cudaArray *d_depth, *d_diffuse, *d_misc, *d_normal, *d_out;
	cudaGraphicsSubResourceGetMappedArray(&d_depth, crc.depth, 0, 0);
	cudaGraphicsSubResourceGetMappedArray(&d_diffuse, crc.diffuse, 0, 0);
	cudaGraphicsSubResourceGetMappedArray(&d_misc, crc.misc, 0, 0);
	cudaGraphicsSubResourceGetMappedArray(&d_normal, crc.normal, 0, 0);
	cudaGraphicsSubResourceGetMappedArray(&d_out, crc.out, 0, 0);

	// Bind the texture references to the cuda arrays
	// Assume channel desc stays the same for the first 4 textures
	cudaChannelFormatDesc channelDesc;
	cudaGetChannelDesc(&channelDesc, d_depth);
	cudaBindTextureToArray(&depthTexRef, d_depth, &channelDesc);
	cudaBindTextureToArray(&diffuseTexRef, d_diffuse, &channelDesc);
	cudaBindTextureToArray(&miscTexRef, d_misc, &channelDesc);
	cudaBindTextureToArray(&normalTexRef, d_normal, &channelDesc);

	// Geez....FINALLY run the kernel
	if(_use41) {
		cudaCalculation41<<<grid, block>>>(pld.numPointLights, (float*)pld.d_pointLightFloats,
			                               (int*)pld.d_pointLightScissorRects, width, height, d_outData);
	}
	else {
		cudaCalculation<<<grid, block>>>(pld.numPointLights, (float*)pld.d_pointLightFloats,
										 (int*)pld.d_pointLightScissorRects, width, height, d_outData);
	}
	
	// Copy the data into the output textures
	cudaMemcpyToArray(d_out, 0, 0, d_outData, sizeof(uchar4)*width*height, cudaMemcpyDeviceToDevice);

	// Unbind the textures from the arrays - just assume there won't be any errors
	cudaUnbindTexture(&depthTexRef);
	cudaUnbindTexture(&diffuseTexRef);
	cudaUnbindTexture(&miscTexRef);
	cudaUnbindTexture(&normalTexRef);
	cudaUnbindTexture(&outTexRef);

	// Unmap the textures from CUDA
	cudaGraphicsUnmapResources(1, &(crc.depth));
	cudaGraphicsUnmapResources(1, &(crc.diffuse));
	cudaGraphicsUnmapResources(1, &(crc.misc));
	cudaGraphicsUnmapResources(1, &(crc.normal));
	cudaGraphicsUnmapResources(1, &(crc.out));


}

int getCUDAShaderIDs(unsigned int **nt, int **sid, dim3 block, int width, int height, bool aFrame, float n, float f) {
if(!_cudaInit) {
		printf("The textures in cuda have not been registered. Cannot perform CUDA calculations.\n");
		*nt = NULL;
		*sid = NULL;
		return 0;
	}
	else if((block.x * block.y) % 2 != 0) {
		printf("Threads per block must be a multiple of 2.\n");
		*nt = NULL;
		*sid = NULL;
		return 0;
	}
	else if((block.x * block.y) < 16) {
		printf("Must have at least 16 threads per block.\n");
		*nt = NULL;
		*sid = NULL;
		return 0;
	}

	dim3 grid = dim3(width / block.x, height / block.y);
	if(width % block.x != 0)
		grid.x = grid.x + 1;
	if(height % block.y != 0)
		grid.y = grid.y + 1;

	if(!(lastBlock.x == block.x && lastBlock.y == block.y && width == lastWidth && height == lastHeight)) {
		delete numTiles;
		numTiles = NULL;
		delete shaderIDs;
		shaderIDs = NULL;
		cudaFree(d_shaderIDs);
		d_shaderIDs = NULL;
		lastBlock = block;
		lastWidth = width;
		lastHeight = height;

		cudaFree(d_outData);
		d_outData = NULL;
	}

	if(!numTiles) {
		numTiles = new unsigned int;
		*numTiles = grid.x * grid.y;
	}
	if(!shaderIDs) {
		shaderIDs = new int[(*numTiles)];
		cudaMalloc(&d_shaderIDs, sizeof(int) * (*numTiles));
	}

	// Get the pointers to the CUDA graphics resources
	cuda_resource_container crc = getCudaResources(aFrame);

	// Map the textures to CUDA
	cudaGraphicsMapResources(1, &(crc.depth));
	cudaGraphicsMapResources(1, &(crc.misc));
	cudaGraphicsMapResources(1, &(crc.miscB));
	cudaGraphicsMapResources(1, &(crc.edge));

	// Get device pointers to the textures
	cudaArray *d_depth, *d_misc, *d_miscB, *d_edge;
	cudaGraphicsSubResourceGetMappedArray(&d_depth, crc.depth, 0, 0);
	cudaGraphicsSubResourceGetMappedArray(&d_misc, crc.misc, 0, 0);
	cudaGraphicsSubResourceGetMappedArray(&d_miscB, crc.miscB, 0, 0);
	cudaGraphicsSubResourceGetMappedArray(&d_edge, crc.edge, 0, 0);

	// Bind the texture references to the cuda arrays
	// Assume channel desc stays the same for the first 4 textures
	cudaChannelFormatDesc channelDesc;
	cudaGetChannelDesc(&channelDesc, d_depth);
	cudaBindTextureToArray(&depthTexRef, d_depth, &channelDesc);
	cudaBindTextureToArray(&miscTexRef, d_misc, &channelDesc);
	cudaBindTextureToArray(&miscBTexRef, d_miscB, &channelDesc);
	cudaBindTextureToArray(&edgeTexRef, d_edge, &channelDesc);

	// Run the shaderID kernel
	if(_use41) {
		generateShaderIDs41<<<grid, block>>>(d_shaderIDs, width, height, n, f);
	}
	else
		generateShaderIDs<<<grid, block>>>(d_shaderIDs, width, height, n, f);

	// Unbind the textures from the arrays - just assume there won't be any errors
	cudaUnbindTexture(&depthTexRef);
	cudaUnbindTexture(&miscTexRef);
	cudaUnbindTexture(&miscBTexRef);
	cudaUnbindTexture(&edgeTexRef);

	// Unmap the textures from CUDA
	cudaGraphicsUnmapResources(1, &(crc.depth));
	cudaGraphicsUnmapResources(1, &(crc.misc));
	cudaGraphicsUnmapResources(1, &(crc.miscB));
	cudaGraphicsUnmapResources(1, &(crc.edge));

	cudaMemcpy(shaderIDs, d_shaderIDs, sizeof(int)*(*numTiles),cudaMemcpyDeviceToHost);
	*nt = numTiles;
	*sid = shaderIDs;

	return grid.x;
}